
#pragma OPENCL EXTENSION cl_khr_fp64 : enable


__kernel void gauss1(__global double *a, const int nn, const int mm, const int n)
{

    //const float epsilon = 1.19209e-07;
    const double epsilon = 2.22045e-16;
    //      const int n = get_global_id(0);


    if (fabs(a[n]) <= epsilon)
        return;

    const int l = get_global_id(0);
    //const int k = get_global_id(1);
    
    const int mb = min(mm, nn - n);

    const int MWMAX = 440;
    const int NWMAX = 440;



    const int M = MWMAX*NWMAX;

    //__private float c;
    int k;

    //const int k = get_global_id(1);

    //__local double a_r[64];        // max mb
    //if (l < mb)
    {
  //      a_r[l] = a[n+M*(l)];
    }

    //barrier(CLK_LOCAL_MEM_FENCE);


    if (l < mb-1)
    {
        __private double c = a[n + M * (l+1)] / a[n];
        //__private double c = native_divide(a[n + M * (l+1)], a[n]);
        //c = a_r[l+1] / a_r[0];
        //__private double c = native_divide(a_r[l+1], a_r[0]);

        if (fabs(c) > epsilon)
        {

            for (k = 0; k < mb-l-1; ++k)
            // if (k < mb-l-1)
            {
//                a[n + M * k + (l+1)] -= c * a_r[k+l+1];
                a[n + M * k + (l+1)] -= c * a[n + M*(k+l+1)];
            }

            a[n + M * (l+1)] = c;
        }

    }

}

